﻿// OCLConv.cpp : 此文件包含 "main" 函数。程序执行将在此处开始并结束。
//

#include "Util.h"


const char* _kernel_convlution_naive = KERNEL(
	__kernel void convlution_naive(int imgOutX, int imgOutY, global float* imgIn, int filterWidth,
		const global float* filter, global float* imgOut) {
	int y = get_global_id(1);
	int x = get_global_id(0);

	int imgInX = imgOutX + filterWidth - 1;//总列数，即每行的宽度
	
	if (y < imgOutY && x < imgOutX) {
		float sum = 0.0;
		for (int fy = 0; fy < filterWidth; fy++) {
			for (int fx = 0; fx < filterWidth; fx++) {
				sum += filter[fy*filterWidth + fx] * imgIn[(x + fx) + (y + fy)*imgInX];
			}
		}
		imgOut[x + y * imgOutX] = sum;
	}
}
);


const char* _kernel_convlution_padding0 = KERNEL(
	__kernel void convlution_padding0(int imgOutX, int imgOutY, global float* imgIn, int filterWidth,
		const global float* filter, global float* imgOut) {
	int y = get_global_id(1);
	int x = get_global_id(0);

	int halfFilterWidth = filterWidth / 2;

	if (y < imgOutY && x < imgOutX) {
		float sum = 0.0;
		for (int fy = -halfFilterWidth; fy <= halfFilterWidth; fy++) {
			for (int fx = -halfFilterWidth; fx <= halfFilterWidth; fx++) {
				if (y + fy >= 0 && y + fy < imgOutY && x + fx >= 0 && x + fx < imgOutX) {
					sum += filter[(fy + halfFilterWidth)*filterWidth + (fx + halfFilterWidth)] * imgIn[(x + fx) + (y + fy)*imgOutX];
				}
			}
		}
		imgOut[x + y * imgOutX] = sum;
	}
}
);



bool check1(int filterWidth, int rowIn, int colIn, int rowOut, int colOut, float* inImg, float* filter, float* outImg) {
	bool correct = 1;
	int padding = (filterWidth - 1) / 2, halfFilterWidth = filterWidth / 2;
	for (int r = 0; r < rowIn && correct; r++)
	{
		for (int c = 0; c < colIn && correct; c++)
		{
			float tempSum = 0.0f;
			for (int i = -halfFilterWidth; i <= halfFilterWidth; i++)
			{
				for (int j = -halfFilterWidth; j <= halfFilterWidth; j++) {
					if (r + i >= 0 && r + i < rowIn&&c + j >= 0 && c + j < colIn) {
						tempSum += inImg[(r + i)*colIn + c + j] * filter[(i + halfFilterWidth) * filterWidth + j + halfFilterWidth];
					}
				}
			}
			if (r - padding >= 0 && r - padding < rowOut&&c - padding >= 0 && c - padding < colOut) {
				if (!floatEq(tempSum, outImg[(r - padding)*colOut + c - padding]))
				{
					printf("Error at [%d, %d], calculation: %f, reference: %f\n", r - padding, c - padding, outImg[(r - padding)*colOut + c - padding], tempSum);
					correct = 0;
				}
			}
		}
	}
	return correct;
}

bool check2(int filterWidth, int rowIn, int colIn, float* inImg, float* filter, float* outImg) {
	bool correct = 1;
	int halfFilterWidth = filterWidth / 2;
	for (int r = 0; r < rowIn && correct; r++)
	{
		for (int c = 0; c < colIn && correct; c++)
		{
			float tempSum = 0.0f;
			for (int i = -halfFilterWidth; i <= halfFilterWidth; i++)
			{
				for (int j = -halfFilterWidth; j <= halfFilterWidth; j++) {
					if (r + i >= 0 && r + i < rowIn&&c + j >= 0 && c + j < colIn) {
						tempSum += inImg[(r + i)*colIn + c + j] * filter[(i + halfFilterWidth) * filterWidth + j + halfFilterWidth];
					}
				}
			}
	
			if (!floatEq(tempSum, outImg[r * colIn + c]))
			{
				printf("Error at [%d, %d], calculation: %f, reference: %f\n", r, c, outImg[r * colIn + c], tempSum);
				correct = 0;
			}
		
		}
	}
	return correct;
}

int conv()
{
	int rowIn = 1024, colIn = 512;
	float* inImg = (float*)malloc(sizeof(float) * rowIn * colIn);
	srand(2);
	for (int i = 0; i < rowIn * colIn; inImg[i] = rand() & 0xF, i++);
    std::cout << "inImg is generated!\n";

	const int filterWidth = 3, filterSize = filterWidth * filterWidth, halfFilterWidth = filterWidth / 2;
	float filter[9] = {
		1.0/9,1.0/9,1.0/9,
		1.0/9,1.0/9,1.0/9,
		1.0/9,1.0/9,1.0/9
	};

	int colOut = 0, rowOut = 0;
	printf("Choose the way of multiply.\n");
	int sel = 0;
	scanf_s("%d", &sel, 1);
	switch (sel)
	{
	case 1:
		colOut = colIn + 1 - filterWidth;
		rowOut = rowIn + 1 - filterWidth;
		break;

	case 2:
		colOut = colIn;
		rowOut = rowIn;
		break;

	default:
		break;
	}
	
	float* outImg = (float*)malloc(sizeof(float)*rowOut*colOut);

	cl_int status;

	OpenCLKernel* m_pOCL = new OpenCLKernel;
	if (0 != m_pOCL->OpenCLInitialize())
	{
		delete m_pOCL;
		m_pOCL = nullptr;
		return EXIT_FAILURE;
	}

	cl_command_queue queue = clCreateCommandQueue(m_pOCL->GetContext(), m_pOCL->GetDevices()[0], NULL, &status);
	GET_CL_ERROR(status, m_pOCL);

	//编译
	size_t nSourceSize = NULL;
	cl_program program = NULL;
	cl_kernel kernel = NULL;
	
	switch (sel)
	{
	case 1:
		nSourceSize = { strlen(_kernel_convlution_naive) };
		program = clCreateProgramWithSource(m_pOCL->GetContext(), 1, &_kernel_convlution_naive, &nSourceSize, &status);
		GET_CL_ERROR(status, m_pOCL);
		status = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
		GET_CL_ERROR(status, m_pOCL);
		kernel = clCreateKernel(program, "convlution_naive", &status);
		GET_CL_ERROR(status, m_pOCL);
		break;
	case 2:
		nSourceSize = { strlen(_kernel_convlution_padding0) };
		program = clCreateProgramWithSource(m_pOCL->GetContext(), 1, &_kernel_convlution_padding0, &nSourceSize, &status);
		GET_CL_ERROR(status, m_pOCL);
		status = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
		GET_CL_ERROR(status, m_pOCL);
		kernel = clCreateKernel(program, "convlution_padding0", &status);
		GET_CL_ERROR(status, m_pOCL);
		break;
	//case 3:
	//	nSourceSize = { strlen(_kernel_multiply_03_) };
	//	program = clCreateProgramWithSource(m_pOCL->GetContext(), 1, &_kernel_multiply_03_, &nSourceSize, &status);
	//	GET_CL_ERROR(status, m_pOCL);
	//	status = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
	//	GET_CL_ERROR(status, m_pOCL);
	//	kernel = clCreateKernel(program, "multiply_03", &status);
	//	GET_CL_ERROR(status, m_pOCL);
	//	break;
	//case 4:
	//	nSourceSize = { strlen(_kernel_multiply_04_) };
	//	program = clCreateProgramWithSource(m_pOCL->GetContext(), 1, &_kernel_multiply_04_, &nSourceSize, &status);
	//	GET_CL_ERROR(status, m_pOCL);
	//	status = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
	//	GET_CL_ERROR(status, m_pOCL);
	//	kernel = clCreateKernel(program, "multiply_04", &status);
	//	GET_CL_ERROR(status, m_pOCL);
	//	break;
	case 0:
	case -1:
	default:
		return EXIT_FAILURE;
	}

	//读内存
	cl_mem bufferIn = clCreateBuffer(m_pOCL->GetContext(), CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float) * rowIn * colIn, inImg, &status);
	GET_CL_ERROR(status, m_pOCL);
	cl_mem bufferFilter = clCreateBuffer(m_pOCL->GetContext(), CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float) * filterSize, filter, &status);
	GET_CL_ERROR(status, m_pOCL);
	cl_mem bufferOut = clCreateBuffer(m_pOCL->GetContext(), CL_MEM_WRITE_ONLY, sizeof(float) * rowOut * colOut, NULL, &status);
	GET_CL_ERROR(status, m_pOCL);
	
	clock_t time_start = clock();

	//参数赋值
	status = clSetKernelArg(kernel, 0, sizeof(int), &colOut);
	GET_CL_ERROR(status, m_pOCL);
	status = clSetKernelArg(kernel, 1, sizeof(int), &rowOut);
	GET_CL_ERROR(status, m_pOCL);
	status = clSetKernelArg(kernel, 2, sizeof(cl_mem), &bufferIn);
	GET_CL_ERROR(status, m_pOCL);
	status = clSetKernelArg(kernel, 3, sizeof(int), &filterWidth);
	GET_CL_ERROR(status, m_pOCL);
	status = clSetKernelArg(kernel, 4, sizeof(cl_mem), &bufferFilter);
	GET_CL_ERROR(status, m_pOCL);
	status = clSetKernelArg(kernel, 5, sizeof(cl_mem), &bufferOut);
	GET_CL_ERROR(status, m_pOCL);


	// 注意不同函数需要调整工作组网格参数
	size_t * globalSize = NULL;
	size_t * localSize = NULL;
	cl_uint work_dim = 0;
	switch (sel)
	{
	case 1:
	case 2:
		work_dim = 2;
		globalSize = new size_t[2];
		globalSize[0] = colOut;
		globalSize[1] = rowOut;
		localSize = new size_t[2];
		localSize[0] = 16;
		localSize[1] = 16;
		break;
	//case 3:
	//	work_dim = 2;
	//	globalSize = new size_t[2];
	//	globalSize[0] = rowA;
	//	globalSize[1] = colB;
	//	localSize = new size_t[2];
	//	localSize[0] = TILE_DIM;
	//	localSize[1] = TILE_DIM;
	//	break;
	//case 2:
	//	work_dim = 1;
	//	globalSize = new size_t[1];
	//	localSize = new size_t[1];
	//	globalSize[0] = rowA;
	//	localSize[0] = 256;
	//	break;
	//case 4: //此算法的size皆以float4为单位
	//	work_dim = 2;
	//	globalSize = new size_t[2];
	//	globalSize[0] = rowA / 4;
	//	globalSize[1] = colB / 4;
	//	localSize = new size_t[2];
	//	localSize[0] = TILE_DIM;
	//	localSize[1] = TILE_DIM;
	//	break;
	default:
		return EXIT_FAILURE;
	}

	//执行
	status = clEnqueueNDRangeKernel(queue, kernel, work_dim, NULL, globalSize, localSize, 0, NULL, NULL);
	GET_CL_ERROR(status, m_pOCL);
	status = clFinish(queue);
	GET_CL_ERROR(status, m_pOCL);

	clock_t time_end = clock();
	printf("\nTime kernel : %d ms\n", time_end - time_start);

	//写内存
	status = clEnqueueReadBuffer(queue, bufferOut, CL_TRUE, 0, sizeof(float) * rowOut * colOut, outImg, 0, NULL, NULL);
	GET_CL_ERROR(status, m_pOCL);


	// 返回并检查结果
	printf("Checking...\n");
	bool correct = 1;
	switch (sel)
	{
	case1:
		correct = check1(filterWidth, rowIn, colIn, rowOut, colOut, inImg, filter, outImg);
		break;

	case 2:
		correct = check2(filterWidth, rowIn, colIn, inImg, filter, outImg);
		break;
	default:
		break;
	}


	if (correct)
		printf("Result correct.\n");

	//结束并释放
	if (globalSize != NULL)
		delete[] globalSize;
	if (localSize != NULL)
		delete[] localSize;
	clReleaseMemObject(bufferIn);
	clReleaseMemObject(bufferFilter);
	clReleaseMemObject(bufferOut);
	clReleaseKernel(kernel);
	clReleaseProgram(program);
	clReleaseCommandQueue(queue);
	delete m_pOCL;
	free(inImg);
	free(outImg);
	
	return 0;
}

// 运行程序: Ctrl + F5 或调试 >“开始执行(不调试)”菜单
// 调试程序: F5 或调试 >“开始调试”菜单

// 入门使用技巧: 
//   1. 使用解决方案资源管理器窗口添加/管理文件
//   2. 使用团队资源管理器窗口连接到源代码管理
//   3. 使用输出窗口查看生成输出和其他消息
//   4. 使用错误列表窗口查看错误
//   5. 转到“项目”>“添加新项”以创建新的代码文件，或转到“项目”>“添加现有项”以将现有代码文件添加到项目
//   6. 将来，若要再次打开此项目，请转到“文件”>“打开”>“项目”并选择 .sln 文件
